#include <iostream>

__global__ void kernel (int x)
{
    asm(".reg .u32 t1;\n\t" // temp reg t1
        " mul.lo.u32 t1, %1, %1;\n\t" // t1 = x * x
        " mul.lo.u32 %0, t1, %1;" // y = t1 * x
        : "=r"(x) : "r" (x));
    printf("x = %d\n", x);
}

__global__ void clz_kernel (int x)
{
    uint32_t count = 0;
    asm("clz.b32 %0, %1;\n\t"
        : "=r"(count) : "r" (x));
     printf("clz count = %u\n", count);
}

__global__ void add_kernel(int x, int y)
{
    int out = 0;
    asm("add.s32 %0, %1, %2;\r\n"
        : "=r"(out) : "r"(x), "r"(y)
    );
    printf("add x = %d\n", out);
}

__global__ void bfind_kernel(int x)
{
    uint32_t out = 0;
    asm("bfind.u32 %0, %1;\r\n"
        : "=r"(out) : "r"(x)
    );
    printf("bfind x = %u\n", out);
}

__global__ void fns_kernel(uint32_t mask, int base, int offset)
{
    uint32_t out = 0;
    asm("fns.b32 %0, %1, %2, %3;\r\n"
        : "=r"(out) : "r"(mask), "r"(base), "r"(offset)
    );
    printf("fns x = %u\n", out);
}

__global__ void bfe_kernel(uint32_t a, uint32_t start, uint32_t len)
{
    uint32_t out = 0;
    asm("bfe.u32 %0, %1, %2, %3;\r\n"
        : "=r"(out) : "r"(a), "r"(start), "r"(len)
    );
    printf("bfe x = %u\n", out);
}

__global__ void bmsk_kernel(uint32_t a, uint32_t len)
{
    uint32_t out = 0;
    asm(".target sm_80\r\n"
        "bmsk.clamp.b32 %0, %1, %2;\r\n"
        : "=r"(out) : "r"(a), "r"(len)
    );
    printf("bmak x = %x\n", out);
}

__global__ void copysign_kernel(float x, float y)
{
    float out = 0;
    asm("copysign.f32 %0, %1, %2;\r\n"
        : "=f"(out) : "f"(x), "f"(y)
    );
    printf("copysign x = %f\n", out);
}

int main() {
    // scanf("%d", &mask);
    // scanf("%d", &base);
    // scanf("%d", &offset);

    // int x = 1;
    // uint32_t count = 88;
    // add_kernel<<<1, 1>>>(x, count);
    // // kernel<<<1, 1>>>(x);
    // clz_kernel<<<1, 1>>>(x);
    // bfind_kernel<<<1, 1>>>(x);

    fns_kernel<<<1,1>>>(0b11001010100110000100001000010000, 5, 2);

    bfe_kernel<<<1,1>>>(0b11001010100110000100001000010000, 0, 5);

    bmsk_kernel<<<1,1>>>(3, 8);

    copysign_kernel<<<1,1>>>(-1.0, 2.0);

    cudaDeviceSynchronize();

    std::cout << "PTX code embedded successfully!" << std::endl;

    return 0;
}